__kernel void matchingObserve(
	__global const float* observeDepthMap,
	__global const float* renderDepthMap, 
	__global const float* observeSkin,
	__global float* sum_1,
	__global float* sum_2,
	__global float* sum_3)
{
	int num = get_global_id(0);
	float sub = fabs(observeDepthMap[num]-renderDepthMap[num]);
	float rMatch;
	if(sub > 0.04f )
	{
		sum_1[num]=4;
	}else{
		
		sum_1[num]=sub*100;
	}

	if(sub < 0.01f && renderDepthMap[num] < 1.f){
		rMatch=1.0f;
	}else{
		rMatch=0.0f;
	}
	sum_2[num] = observeSkin[num] != 1.f || rMatch == 1.f;
	sum_3[num] = observeSkin[num] != 1.f && rMatch == 1.f;
}

__kernel void reduction_vector(
	__global float4* data, 
    __local float4* partial_sums,
	__global float* output) {

   int lid = get_local_id(0);
   int group_size = get_local_size(0);

   partial_sums[lid] = data[get_global_id(0)];
   barrier(CLK_LOCAL_MEM_FENCE);

   for(int i = group_size/2; i>0; i >>= 1) {
      if(lid < i) {
         partial_sums[lid] += partial_sums[lid + i];
      }
      barrier(CLK_LOCAL_MEM_FENCE);
   }

   if(lid == 0) {
      output[get_group_id(0)] = dot(partial_sums[0], (float4)(1.0f));
   }
}